#include <fstream>
#include <omdcl/machine.hpp>
#include <omdcl/program.hpp>

using std::ifstream;

void MDComputeProgram::ReadProgramFile(string source) {
	assert(file_exist(source), "can not find device program source ("+source+")");	
	ifstream fsrc;
	fsrc.open(source.c_str());
	assert(fsrc.good(), "bad file: "+source);
	
	string sbuff;
	char sline[2048]; // max 2047 characters per line!!
	
	while(fsrc.good()) {
        memset(sline,0,2048);
		fsrc.getline(sline,2048);
		if(fsrc.good()) {
		    // reserved for a special commands...
		    if(sline[0]!='#'&&sline[1]!='!') sbuff.append(sline);
		    sbuff.append("\n");
        }
	}
	
	fsrc.close();
	source_code=new char[sbuff.size()+1];
	memset(source_code,0,sbuff.size()+1);
	sbuff.copy(source_code, sbuff.size());
}

void MDComputeProgram::Create() {
	cl_int err;
	program=clCreateProgramWithSource(machine->context,1,(const char**)&source_code,NULL,&err);
	assert(err==CL_SUCCESS, "failed creating program");
	
	err = clBuildProgram(program, 0, NULL, compile_options, NULL, NULL);

	if (err != CL_SUCCESS) {
		size_t len;
		char buffer[2048];
		string serr("Error: Failed to build program executable!\n\n");
		clGetProgramBuildInfo(program, machine->device, CL_PROGRAM_BUILD_LOG,
							  sizeof(buffer), buffer, &len);
		serr.append(buffer);
		serr.append("\nsource:\n");
		serr.append(source_code);
		die(serr);
	}
	
	cl_uint numof_kernels;
	cl_kernel kers[OMDCL_MAX_KERNEL];
	err=clCreateKernelsInProgram(program,OMDCL_MAX_KERNEL,kers,&numof_kernels);
	assert(err==CL_SUCCESS, "failed creating kernels from program");

	for(cl_uint i=0;i<numof_kernels;i++) {
		char st[128];
		memset(st,0,128);
		clGetKernelInfo(kers[i],CL_KERNEL_FUNCTION_NAME,127,st,NULL);
		KernelPack* krn=new KernelPack(st,kers[i]);
		if(krn->name=="clmain")active_kernel=krn;
		kernels.push_back(krn);
	}
	
	nd_dimension=0;
	memset(nd_local,0,3*sizeof(size_t));
	memset(nd_global,0,3*sizeof(size_t));
	memset(nd_offset,0,3*sizeof(size_t));	
}

void MDComputeProgram::SetDimension(size_t cx) {
	// one dimensional work-group and work items...
	nd_dimension=1;
	nd_local[1]=nd_local[2]=0;
	nd_global[1]=nd_global[2]=0;
	
	assert(CL_SUCCESS==clGetKernelWorkGroupInfo(
	                    active_kernel->kernel,
	                    machine->device, 
	                    CL_KERNEL_WORK_GROUP_SIZE,
                        sizeof(size_t),nd_local, NULL),
            			"failed to get workgroup info");
        
    size_t rem=cx%nd_local[0];
    if(rem!=0)  {
		nd_global[0]=(nd_local[0])*(1+cx/nd_local[0]);
	} else nd_global[0]=cx;  
}

void MDComputeProgram::ClearArgument() {
	assert(active_kernel, "no active kernel selected");
	active_kernel->argcount=0;
}

void MDComputeProgram::PushArgument(size_t size, const void* arg) {
	assert(active_kernel, "no active kernel selected");
	clSetKernelArg(active_kernel->kernel,active_kernel->argcount++,size,arg);
}

KernelPack* MDComputeProgram::SelectKernel(string kernel_name) {
	for(int i=0;i<(int)kernels.size();i++) {
		if(kernels[i]->name==kernel_name) {
			kernels[i]->argcount=0;
			active_kernel=kernels[i];
			return kernels[i];
		}
	}
	die("kernel ("+kernel_name+") not found");
	return NULL;
}

void MDComputeProgram::Execute() {
	// evt==null means blocking...
	if(!nd_dimension) die("uninitiated ndrange... use SetDimension!");
	if(CL_SUCCESS!=clEnqueueNDRangeKernel(machine->queue,active_kernel->kernel,
										  nd_dimension,NULL,
										  nd_global, nd_local, 0, NULL, NULL))
		die("failed enqueuing program");
}

void MDComputeProgram::Execute(cl_event& evt) {
	// evt==null means blocking...
	if(!nd_dimension) die("uninitiated ndrange... use SetDimension!");
	if(CL_SUCCESS!=clEnqueueNDRangeKernel(machine->queue,active_kernel->kernel,
										  nd_dimension,NULL,
										  nd_global, nd_local, 0, NULL, &evt))
		die("failed enqueuing program");
}

MDComputeProgram::MDComputeProgram(MDComputeMachine* mac, 
								 string programpath, 
								 string cplopt) {
    set_name("omdcl program");
    register_class(get_name());

	source_code=NULL;
	active_kernel=NULL;
	memset(compile_options,0x0,1024);
	cplopt.copy(compile_options,cplopt.size());
	machine=mac;
	ReadProgramFile(programpath);
	Create();
}

MDComputeProgram::~MDComputeProgram() {
	delete[] source_code;
	for(int i=0;i<(int)kernels.size();i++) {
        clReleaseKernel(kernels[i]->kernel);
	}
	clReleaseProgram(program);
}
